#include <stdio.h>
#include <cmath>


__global__ void my_cuda_func();

int main()
{
    cudaDeviceProp my_device_props;
    cudaGetDeviceProperties(&my_device_props, 0);
    auto max_thd_per_blk = my_device_props.maxThreadsPerBlock;
    auto max_thd_x_dim = my_device_props.maxThreadsDim[0];
    auto max_thd_y_dim = my_device_props.maxThreadsDim[1];
    auto max_thd_z_dim = my_device_props.maxThreadsDim[2];

    printf("The maximum number of threads in a block is %d\n", max_thd_per_blk);
    printf("The maximum number of threads in the x dim of a block is %d\n", max_thd_x_dim);
    printf("The maximum number of threads in the y dim of a block is %d\n", max_thd_y_dim);
    printf("The maximum number of threads in the z dim of a block is %d\n", max_thd_z_dim);

    // shared memory of one thread block
    auto max_shared_mem_per_thd_blk = my_device_props.sharedMemPerBlock;
    printf("The maximum shared memory of a thread block is %zu Bytes,\n"
           "and is equal to %u KB as well.\n"
           "What's more, it can hold maximum %u float point number.\n", 
           max_shared_mem_per_thd_blk, (unsigned int)(max_shared_mem_per_thd_blk / pow(2, 10)), 
           (unsigned int)(max_shared_mem_per_thd_blk / 4));
    auto max_shared_mem_per_sm = my_device_props.sharedMemPerMultiprocessor;
    printf("The maximum shared memory of a thread block is %zu Bytes, \n"
           "and is equal to %u KB as well.\n", max_shared_mem_per_sm, 
           (unsigned int)(max_shared_mem_per_sm / pow(2, 10)));

    dim3 grid_layout(2, 3, 3);
    dim3 block_layout(2, 3, 4);
    my_cuda_func<<<grid_layout, block_layout>>>();
    cudaDeviceSynchronize();
}

__global__ void my_cuda_func()
{
    const unsigned thd_local_idx = threadIdx.x + threadIdx.y * blockDim.x + \
                              threadIdx.z * blockDim.y * blockDim.x;
    const unsigned int thds_per_blk = blockDim.z * blockDim.y * blockDim.x;
    const unsigned int blk_idx = blockIdx.x + blockIdx.y * gridDim.x + \
                                 blockIdx.z * gridDim.y * gridDim.x;
    const unsigned int thd_global_idx = thd_local_idx + thds_per_blk * blk_idx;

    printf("The global index of current thread is %u!\n", thd_global_idx);
}
